<!-- HTML header for doxygen 1.8.3.1-->
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.4"/>
<title>CUB: cub::BlockDiscontinuity&lt; T, BLOCK_THREADS &gt; Class Template Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
<link href="extra_stylesheet.css" rel="stylesheet" type="text/css"/>
<link rel="shortcut icon" href="favicon.ico" type="image/x-icon" />
<script type="text/javascript">
  var _gaq = _gaq || [];
  _gaq.push(['_setAccount', 'UA-38890655-1']);
  _gaq.push(['_trackPageview']);
  (function() {
    var ga = document.createElement('script'); ga.type = 'text/javascript'; ga.async = true;
    ga.src = ('https:' == document.location.protocol ? 'https://ssl' : 'http://www') + '.google-analytics.com/ga.js';
    var s = document.getElementsByTagName('script')[0]; s.parentNode.insertBefore(ga, s);
  })();
</script>
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td style="padding-left: 0.5em;">
   <div id="projectname">CUB
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.4 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Main&#160;Page</span></a></li>
      <li><a href="modules.html"><span>Modules</span></a></li>
      <li class="current"><a href="annotated.html"><span>Classes</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
  <div id="navrow2" class="tabs2">
    <ul class="tablist">
      <li><a href="annotated.html"><span>Class&#160;List</span></a></li>
      <li><a href="classes.html"><span>Class&#160;Index</span></a></li>
    </ul>
  </div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Namespaces</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Typedefs</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(7)"><span class="SelectionMark">&#160;</span>Enumerations</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(8)"><span class="SelectionMark">&#160;</span>Enumerator</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(9)"><span class="SelectionMark">&#160;</span>Groups</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div id="nav-path" class="navpath">
  <ul>
<li class="navelem"><a class="el" href="namespacecub.html">cub</a></li><li class="navelem"><a class="el" href="classcub_1_1_block_discontinuity.html">BlockDiscontinuity</a></li>  </ul>
</div>
</div><!-- top -->
<div class="header">
  <div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="classcub_1_1_block_discontinuity-members.html">List of all members</a>  </div>
  <div class="headertitle">
<div class="title">cub::BlockDiscontinuity&lt; T, BLOCK_THREADS &gt; Class Template Reference<div class="ingroups"><a class="el" href="group___block_module.html">Block-wide</a></div></div>  </div>
</div><!--header-->
<div class="contents">
<a name="details" id="details"></a><h2 class="groupheader">Detailed description</h2>
<div class="textblock"><h3>template&lt;
    typename T, 
    int BLOCK_THREADS&gt;<br/>
class cub::BlockDiscontinuity&lt; T, BLOCK_THREADS &gt;</h3>

<p>The <a class="el" href="classcub_1_1_block_discontinuity.html" title="The BlockDiscontinuity class provides collective methods for flagging discontinuities within an order...">BlockDiscontinuity</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for flagging discontinuities within an ordered set of items partitioned across a CUDA thread block. </p>
<div class="image">
<img src="discont_logo.png" alt="discont_logo.png"/>
<div class="caption">
.</div></div>
 <dl class="section user"><dt>Overview</dt><dd>A set of "head flags" (or "tail flags") is often used to indicate corresponding items that differ from their predecessors (or successors). For example, head flags are convenient for demarcating disjoint data segments as part of a segmented scan or reduction.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">T</td><td>The data type to be flagged. </td></tr>
    <tr><td class="paramname">BLOCK_THREADS</td><td>The thread block size in threads.</td></tr>
  </table>
  </dd>
</dl>
<dl class="section user"><dt>A Simple Example</dt><dd>Every thread in the block uses the <a class="el" href="classcub_1_1_block_discontinuity.html" title="The BlockDiscontinuity class provides collective methods for flagging discontinuities within an order...">BlockDiscontinuity</a> class by first specializing the <a class="el" href="classcub_1_1_block_discontinuity.html" title="The BlockDiscontinuity class provides collective methods for flagging discontinuities within an order...">BlockDiscontinuity</a> type, then instantiating an instance with parameters for communication, and finally invoking collective member functions. </dd></dl>
<dl class="section user"><dt></dt><dd>The code snippet below illustrates the head flagging of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </dd></dl>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockDiscontinuity for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_discontinuity.html" title="The BlockDiscontinuity class provides collective methods for flagging discontinuities within an order...">cub::BlockDiscontinuity&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_discontinuity.html#a5bee18d5e423a26fd3fefa741832ad8c" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockDiscontinuity</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockDiscontinuity</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockDiscontinuity::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute head flags for discontinuities in the segment</span></div>
<div class="line">    <span class="keywordtype">int</span> head_flags[4];</div>
<div class="line">    <a class="code" href="classcub_1_1_block_discontinuity.html#a5bee18d5e423a26fd3fefa741832ad8c" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockDiscontinuity</a>(temp_storage).FlagHeads(head_flags, thread_data, <a class="code" href="structcub_1_1_inequality.html" title="Default inequality functor. ">cub::Inequality</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [0,0,1,1], [1,1,1,1], [2,3,3,3], [3,4,4,4], ... }</code>. The corresponding output <code>head_flags</code> in those threads will be <code>{ [1,0,1,0], [0,0,0,0], [1,1,0,0], [0,1,0,0], ... }</code>.</dd></dl>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>Zero bank conflicts for most types. </li>
</ul>
</dd></dl>

<p>Definition at line <a class="el" href="block__discontinuity_8cuh_source.html#l00097">97</a> of file <a class="el" href="block__discontinuity_8cuh_source.html">block_discontinuity.cuh</a>.</p>
</div><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_block_discontinuity_1_1_temp_storage.html">TempStorage</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The operations exposed by <a class="el" href="classcub_1_1_block_discontinuity.html" title="The BlockDiscontinuity class provides collective methods for flagging discontinuities within an order...">BlockDiscontinuity</a> require a temporary memory allocation of this nested type for thread communication. This opaque storage can be allocated directly using the <code>__shared__</code> keyword. Alternatively, it can be aliased to externally allocated memory (shared or global) or <code>union</code>'d with other storage allocation types to facilitate memory reuse.  <a href="structcub_1_1_block_discontinuity_1_1_temp_storage.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="pub-methods"></a>
Public Methods</h2></td></tr>
<tr><td colspan="2"><div class="groupHeader">Collective constructors</div></td></tr>
<tr class="memitem:a5bee18d5e423a26fd3fefa741832ad8c"><td class="memItemLeft" align="right" valign="top"><a class="anchor" id="a5bee18d5e423a26fd3fefa741832ad8c"></a>
__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_discontinuity.html#a5bee18d5e423a26fd3fefa741832ad8c">BlockDiscontinuity</a> ()</td></tr>
<tr class="memdesc:a5bee18d5e423a26fd3fefa741832ad8c"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Threads are identified using <code>threadIdx.x</code>. <br/></td></tr>
<tr class="separator:a5bee18d5e423a26fd3fefa741832ad8c"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a3f5c69462af0c9609f845e474842cfb2"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_discontinuity.html#a3f5c69462af0c9609f845e474842cfb2">BlockDiscontinuity</a> (<a class="el" href="structcub_1_1_block_discontinuity_1_1_temp_storage.html">TempStorage</a> &amp;temp_storage)</td></tr>
<tr class="memdesc:a3f5c69462af0c9609f845e474842cfb2"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <code>threadIdx.x</code>.  <a href="#a3f5c69462af0c9609f845e474842cfb2">More...</a><br/></td></tr>
<tr class="separator:a3f5c69462af0c9609f845e474842cfb2"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a4c99f418ce99535494ae1c309f58f0ac"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_discontinuity.html#a4c99f418ce99535494ae1c309f58f0ac">BlockDiscontinuity</a> (int linear_tid)</td></tr>
<tr class="memdesc:a4c99f418ce99535494ae1c309f58f0ac"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier.  <a href="#a4c99f418ce99535494ae1c309f58f0ac">More...</a><br/></td></tr>
<tr class="separator:a4c99f418ce99535494ae1c309f58f0ac"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a0edc454bf6d30810b3b086ec77ed4e11"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_discontinuity.html#a0edc454bf6d30810b3b086ec77ed4e11">BlockDiscontinuity</a> (<a class="el" href="structcub_1_1_block_discontinuity_1_1_temp_storage.html">TempStorage</a> &amp;temp_storage, int linear_tid)</td></tr>
<tr class="memdesc:a0edc454bf6d30810b3b086ec77ed4e11"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier.  <a href="#a0edc454bf6d30810b3b086ec77ed4e11">More...</a><br/></td></tr>
<tr class="separator:a0edc454bf6d30810b3b086ec77ed4e11"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Head flag operations</div></td></tr>
<tr class="memitem:afabe7f18d7098e7d0f149696bec17186"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD, typename FlagT , typename FlagOp &gt; </td></tr>
<tr class="memitem:afabe7f18d7098e7d0f149696bec17186"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_discontinuity.html#afabe7f18d7098e7d0f149696bec17186">FlagHeads</a> (FlagT(&amp;head_flags)[ITEMS_PER_THREAD], T(&amp;input)[ITEMS_PER_THREAD], FlagOp flag_op)</td></tr>
<tr class="memdesc:afabe7f18d7098e7d0f149696bec17186"><td class="mdescLeft">&#160;</td><td class="mdescRight">Sets head flags indicating discontinuities between items partitioned across the thread block, for which the first item has no reference and is always flagged.  <a href="#afabe7f18d7098e7d0f149696bec17186">More...</a><br/></td></tr>
<tr class="separator:afabe7f18d7098e7d0f149696bec17186"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ac1d38030af4f9228613f78207e64a314"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD, typename FlagT , typename FlagOp &gt; </td></tr>
<tr class="memitem:ac1d38030af4f9228613f78207e64a314"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_discontinuity.html#ac1d38030af4f9228613f78207e64a314">FlagHeads</a> (FlagT(&amp;head_flags)[ITEMS_PER_THREAD], T(&amp;input)[ITEMS_PER_THREAD], FlagOp flag_op, T tile_predecessor_item)</td></tr>
<tr class="memdesc:ac1d38030af4f9228613f78207e64a314"><td class="mdescLeft">&#160;</td><td class="mdescRight">Sets head flags indicating discontinuities between items partitioned across the thread block.  <a href="#ac1d38030af4f9228613f78207e64a314">More...</a><br/></td></tr>
<tr class="separator:ac1d38030af4f9228613f78207e64a314"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Tail flag operations</div></td></tr>
<tr class="memitem:acf705e15934ba04c556e5860a2853652"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD, typename FlagT , typename FlagOp &gt; </td></tr>
<tr class="memitem:acf705e15934ba04c556e5860a2853652"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_discontinuity.html#acf705e15934ba04c556e5860a2853652">FlagTails</a> (FlagT(&amp;tail_flags)[ITEMS_PER_THREAD], T(&amp;input)[ITEMS_PER_THREAD], FlagOp flag_op)</td></tr>
<tr class="memdesc:acf705e15934ba04c556e5860a2853652"><td class="mdescLeft">&#160;</td><td class="mdescRight">Sets tail flags indicating discontinuities between items partitioned across the thread block, for which the last item has no reference and is always flagged.  <a href="#acf705e15934ba04c556e5860a2853652">More...</a><br/></td></tr>
<tr class="separator:acf705e15934ba04c556e5860a2853652"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ad80f6195b0366d7f88180dc349043179"><td class="memTemplParams" colspan="2">template&lt;int ITEMS_PER_THREAD, typename FlagT , typename FlagOp &gt; </td></tr>
<tr class="memitem:ad80f6195b0366d7f88180dc349043179"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_discontinuity.html#ad80f6195b0366d7f88180dc349043179">FlagTails</a> (FlagT(&amp;tail_flags)[ITEMS_PER_THREAD], T(&amp;input)[ITEMS_PER_THREAD], FlagOp flag_op, T tile_successor_item)</td></tr>
<tr class="memdesc:ad80f6195b0366d7f88180dc349043179"><td class="mdescLeft">&#160;</td><td class="mdescRight">Sets tail flags indicating discontinuities between items partitioned across the thread block.  <a href="#ad80f6195b0366d7f88180dc349043179">More...</a><br/></td></tr>
<tr class="separator:ad80f6195b0366d7f88180dc349043179"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
<h2 class="groupheader">Constructor &amp; Destructor Documentation</h2>
<a class="anchor" id="a3f5c69462af0c9609f845e474842cfb2"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_block_discontinuity.html">cub::BlockDiscontinuity</a>&lt; T, BLOCK_THREADS &gt;::<a class="el" href="classcub_1_1_block_discontinuity.html">BlockDiscontinuity</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structcub_1_1_block_discontinuity_1_1_temp_storage.html">TempStorage</a> &amp;&#160;</td>
          <td class="paramname"><em>temp_storage</em>)</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <code>threadIdx.x</code>. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">temp_storage</td><td>Reference to memory allocation having layout type <a class="el" href="structcub_1_1_block_discontinuity_1_1_temp_storage.html" title="The operations exposed by BlockDiscontinuity require a temporary memory allocation of this nested typ...">TempStorage</a> </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__discontinuity_8cuh_source.html#l00179">179</a> of file <a class="el" href="block__discontinuity_8cuh_source.html">block_discontinuity.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a4c99f418ce99535494ae1c309f58f0ac"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_block_discontinuity.html">cub::BlockDiscontinuity</a>&lt; T, BLOCK_THREADS &gt;::<a class="el" href="classcub_1_1_block_discontinuity.html">BlockDiscontinuity</a> </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>)</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__discontinuity_8cuh_source.html#l00190">190</a> of file <a class="el" href="block__discontinuity_8cuh_source.html">block_discontinuity.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a0edc454bf6d30810b3b086ec77ed4e11"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_block_discontinuity.html">cub::BlockDiscontinuity</a>&lt; T, BLOCK_THREADS &gt;::<a class="el" href="classcub_1_1_block_discontinuity.html">BlockDiscontinuity</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structcub_1_1_block_discontinuity_1_1_temp_storage.html">TempStorage</a> &amp;&#160;</td>
          <td class="paramname"><em>temp_storage</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">temp_storage</td><td>Reference to memory allocation having layout type <a class="el" href="structcub_1_1_block_discontinuity_1_1_temp_storage.html" title="The operations exposed by BlockDiscontinuity require a temporary memory allocation of this nested typ...">TempStorage</a> </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td><b>[optional]</b> A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__discontinuity_8cuh_source.html#l00201">201</a> of file <a class="el" href="block__discontinuity_8cuh_source.html">block_discontinuity.cuh</a>.</p>

</div>
</div>
<h2 class="groupheader">Member Function Documentation</h2>
<a class="anchor" id="afabe7f18d7098e7d0f149696bec17186"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD, typename FlagT , typename FlagOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_discontinuity.html">cub::BlockDiscontinuity</a>&lt; T, BLOCK_THREADS &gt;::FlagHeads </td>
          <td>(</td>
          <td class="paramtype">FlagT(&amp;)&#160;</td>
          <td class="paramname"><em>head_flags</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">FlagOp&#160;</td>
          <td class="paramname"><em>flag_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Sets head flags indicating discontinuities between items partitioned across the thread block, for which the first item has no reference and is always flagged. </p>
<p>The flag <code>head_flags<sub><em>i</em></sub></code> is set for item <code>input<sub><em>i</em></sub></code> when <code>flag_op(</code><em>previous-item</em><code>, input<sub><em>i</em></sub>)</code> returns <code>true</code> (where <em>previous-item</em> is either the preceding item in the same thread or the last item in the previous thread). Furthermore, <code>head_flags<sub><em>i</em></sub></code> is always set for <code>input&gt;<sub>0</sub></code> in <em>thread</em><sub>0</sub>.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates the head-flagging of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockDiscontinuity for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_discontinuity.html" title="The BlockDiscontinuity class provides collective methods for flagging discontinuities within an order...">cub::BlockDiscontinuity&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_discontinuity.html#a5bee18d5e423a26fd3fefa741832ad8c" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockDiscontinuity</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockDiscontinuity</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockDiscontinuity::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute head flags for discontinuities in the segment</span></div>
<div class="line">    <span class="keywordtype">int</span> head_flags[4];</div>
<div class="line">    <a class="code" href="classcub_1_1_block_discontinuity.html#a5bee18d5e423a26fd3fefa741832ad8c" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockDiscontinuity</a>(temp_storage).FlagHeads(head_flags, thread_data, <a class="code" href="structcub_1_1_inequality.html" title="Default inequality functor. ">cub::Inequality</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [0,0,1,1], [1,1,1,1], [2,3,3,3], [3,4,4,4], ... }</code>. The corresponding output <code>head_flags</code> in those threads will be <code>{ [1,0,1,0], [0,0,0,0], [1,1,0,0], [0,1,0,0], ... }</code>.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">FlagT</td><td><b>[inferred]</b> The flag type (must be an integer type) </td></tr>
    <tr><td class="paramname">FlagOp</td><td><b>[inferred]</b> Binary predicate functor type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> or member <code>T operator()(const T &amp;a, const T &amp;b, unsigned int b_index)</code>, and returning <code>true</code> if a discontinuity exists between <code>a</code> and <code>b</code>, otherwise <code>false</code>. <code>b_index</code> is the rank of b in the aggregate tile of data. </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[out]</td><td class="paramname">head_flags</td><td>Calling thread's discontinuity head_flags </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">flag_op</td><td>Binary boolean flag predicate </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__discontinuity_8cuh_source.html#l00271">271</a> of file <a class="el" href="block__discontinuity_8cuh_source.html">block_discontinuity.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ac1d38030af4f9228613f78207e64a314"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD, typename FlagT , typename FlagOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_discontinuity.html">cub::BlockDiscontinuity</a>&lt; T, BLOCK_THREADS &gt;::FlagHeads </td>
          <td>(</td>
          <td class="paramtype">FlagT(&amp;)&#160;</td>
          <td class="paramname"><em>head_flags</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">FlagOp&#160;</td>
          <td class="paramname"><em>flag_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>tile_predecessor_item</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Sets head flags indicating discontinuities between items partitioned across the thread block. </p>
<p>The flag <code>head_flags<sub><em>i</em></sub></code> is set for item <code>input<sub><em>i</em></sub></code> when <code>flag_op(</code><em>previous-item</em><code>, input<sub><em>i</em></sub>)</code> returns <code>true</code> (where <em>previous-item</em> is either the preceding item in the same thread or the last item in the previous thread). For <em>thread</em><sub>0</sub>, item <code>input<sub>0</sub></code> is compared against <code>tile_predecessor_item</code>.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates the head-flagging of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockDiscontinuity for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_discontinuity.html" title="The BlockDiscontinuity class provides collective methods for flagging discontinuities within an order...">cub::BlockDiscontinuity&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_discontinuity.html#a5bee18d5e423a26fd3fefa741832ad8c" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockDiscontinuity</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockDiscontinuity</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockDiscontinuity::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have thread0 obtain the predecessor item for the entire tile</span></div>
<div class="line">    <span class="keywordtype">int</span> tile_predecessor_item;</div>
<div class="line">    <span class="keywordflow">if</span> (threadIdx.x == 0) tile_predecessor_item == ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute head flags for discontinuities in the segment</span></div>
<div class="line">    <span class="keywordtype">int</span> head_flags[4];</div>
<div class="line">    <a class="code" href="classcub_1_1_block_discontinuity.html#a5bee18d5e423a26fd3fefa741832ad8c" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockDiscontinuity</a>(temp_storage).FlagHeads(</div>
<div class="line">        head_flags, thread_data, <a class="code" href="structcub_1_1_inequality.html" title="Default inequality functor. ">cub::Inequality</a>(), tile_predecessor_item);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [0,0,1,1], [1,1,1,1], [2,3,3,3], [3,4,4,4], ... }</code>, and that <code>tile_predecessor_item</code> is <code>0</code>. The corresponding output <code>head_flags</code> in those threads will be <code>{ [0,0,1,0], [0,0,0,0], [1,1,0,0], [0,1,0,0], ... }</code>.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">FlagT</td><td><b>[inferred]</b> The flag type (must be an integer type) </td></tr>
    <tr><td class="paramname">FlagOp</td><td><b>[inferred]</b> Binary predicate functor type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> or member <code>T operator()(const T &amp;a, const T &amp;b, unsigned int b_index)</code>, and returning <code>true</code> if a discontinuity exists between <code>a</code> and <code>b</code>, otherwise <code>false</code>. <code>b_index</code> is the rank of b in the aggregate tile of data. </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[out]</td><td class="paramname">head_flags</td><td>Calling thread's discontinuity head_flags </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">flag_op</td><td>Binary boolean flag predicate </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">tile_predecessor_item</td><td><b>[<em>thread</em><sub>0</sub> only]</b> Item with which to compare the first tile item (<code>input<sub>0</sub></code> from <em>thread</em><sub>0</sub>). </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__discontinuity_8cuh_source.html#l00361">361</a> of file <a class="el" href="block__discontinuity_8cuh_source.html">block_discontinuity.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="acf705e15934ba04c556e5860a2853652"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD, typename FlagT , typename FlagOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_discontinuity.html">cub::BlockDiscontinuity</a>&lt; T, BLOCK_THREADS &gt;::FlagTails </td>
          <td>(</td>
          <td class="paramtype">FlagT(&amp;)&#160;</td>
          <td class="paramname"><em>tail_flags</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">FlagOp&#160;</td>
          <td class="paramname"><em>flag_op</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Sets tail flags indicating discontinuities between items partitioned across the thread block, for which the last item has no reference and is always flagged. </p>
<p>The flag <code>tail_flags<sub><em>i</em></sub></code> is set for item <code>input<sub><em>i</em></sub></code> when <code>flag_op(input<sub><em>i</em></sub>, </code><em>next-item</em><code>)</code> returns <code>true</code> (where <em>next-item</em> is either the next item in the same thread or the first item in the next thread). Furthermore, <code>tail_flags<sub>ITEMS_PER_THREAD-1</sub></code> is always set for <em>thread</em><sub><code>BLOCK_THREADS</code>-1</sub>.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates the tail-flagging of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockDiscontinuity for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_discontinuity.html" title="The BlockDiscontinuity class provides collective methods for flagging discontinuities within an order...">cub::BlockDiscontinuity&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_discontinuity.html#a5bee18d5e423a26fd3fefa741832ad8c" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockDiscontinuity</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockDiscontinuity</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockDiscontinuity::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute tail flags for discontinuities in the segment</span></div>
<div class="line">    <span class="keywordtype">int</span> tail_flags[4];</div>
<div class="line">    <a class="code" href="classcub_1_1_block_discontinuity.html#a5bee18d5e423a26fd3fefa741832ad8c" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockDiscontinuity</a>(temp_storage).FlagTails(tail_flags, thread_data, <a class="code" href="structcub_1_1_inequality.html" title="Default inequality functor. ">cub::Inequality</a>());</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [0,0,1,1], [1,1,1,1], [2,3,3,3], ..., [124,125,125,125] }</code>. The corresponding output <code>tail_flags</code> in those threads will be <code>{ [0,1,0,0], [0,0,0,1], [1,0,0,...], ..., [1,0,0,1] }</code>.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">FlagT</td><td><b>[inferred]</b> The flag type (must be an integer type) </td></tr>
    <tr><td class="paramname">FlagOp</td><td><b>[inferred]</b> Binary predicate functor type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> or member <code>T operator()(const T &amp;a, const T &amp;b, unsigned int b_index)</code>, and returning <code>true</code> if a discontinuity exists between <code>a</code> and <code>b</code>, otherwise <code>false</code>. <code>b_index</code> is the rank of b in the aggregate tile of data. </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[out]</td><td class="paramname">tail_flags</td><td>Calling thread's discontinuity tail_flags </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">flag_op</td><td>Binary boolean flag predicate </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__discontinuity_8cuh_source.html#l00456">456</a> of file <a class="el" href="block__discontinuity_8cuh_source.html">block_discontinuity.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ad80f6195b0366d7f88180dc349043179"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS&gt; </div>
<div class="memtemplate">
template&lt;int ITEMS_PER_THREAD, typename FlagT , typename FlagOp &gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_discontinuity.html">cub::BlockDiscontinuity</a>&lt; T, BLOCK_THREADS &gt;::FlagTails </td>
          <td>(</td>
          <td class="paramtype">FlagT(&amp;)&#160;</td>
          <td class="paramname"><em>tail_flags</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>input</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">FlagOp&#160;</td>
          <td class="paramname"><em>flag_op</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>tile_successor_item</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Sets tail flags indicating discontinuities between items partitioned across the thread block. </p>
<p>The flag <code>tail_flags<sub><em>i</em></sub></code> is set for item <code>input<sub><em>i</em></sub></code> when <code>flag_op(input<sub><em>i</em></sub>, </code><em>next-item</em><code>)</code> returns <code>true</code> (where <em>next-item</em> is either the next item in the same thread or the first item in the next thread). For <em>thread</em><sub><em>BLOCK_THREADS</em>-1</sub>, item <code>input</code><sub><em>ITEMS_PER_THREAD</em>-1</sub> is compared against <code>tile_predecessor_item</code>.</p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates the tail-flagging of 512 integer items that are partitioned in a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across 128 threads where each thread owns 4 consecutive items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockDiscontinuity for 128 threads on type int</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_discontinuity.html" title="The BlockDiscontinuity class provides collective methods for flagging discontinuities within an order...">cub::BlockDiscontinuity&lt;int, 128&gt;</a> <a class="code" href="classcub_1_1_block_discontinuity.html#a5bee18d5e423a26fd3fefa741832ad8c" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockDiscontinuity</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockDiscontinuity</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockDiscontinuity::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Have thread127 obtain the successor item for the entire tile</span></div>
<div class="line">    <span class="keywordtype">int</span> tile_successor_item;</div>
<div class="line">    <span class="keywordflow">if</span> (threadIdx.x == 127) tile_successor_item == ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively compute tail flags for discontinuities in the segment</span></div>
<div class="line">    <span class="keywordtype">int</span> tail_flags[4];</div>
<div class="line">    <a class="code" href="classcub_1_1_block_discontinuity.html#a5bee18d5e423a26fd3fefa741832ad8c" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockDiscontinuity</a>(temp_storage).FlagTails(</div>
<div class="line">        tail_flags, thread_data, <a class="code" href="structcub_1_1_inequality.html" title="Default inequality functor. ">cub::Inequality</a>(), tile_successor_item);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of input <code>thread_data</code> across the block of threads is <code>{ [0,0,1,1], [1,1,1,1], [2,3,3,3], ..., [124,125,125,125] }</code> and that <code>tile_successor_item</code> is <code>125</code>. The corresponding output <code>tail_flags</code> in those threads will be <code>{ [0,1,0,0], [0,0,0,1], [1,0,0,...], ..., [1,0,0,0] }</code>.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">FlagT</td><td><b>[inferred]</b> The flag type (must be an integer type) </td></tr>
    <tr><td class="paramname">FlagOp</td><td><b>[inferred]</b> Binary predicate functor type having member <code>T operator()(const T &amp;a, const T &amp;b)</code> or member <code>T operator()(const T &amp;a, const T &amp;b, unsigned int b_index)</code>, and returning <code>true</code> if a discontinuity exists between <code>a</code> and <code>b</code>, otherwise <code>false</code>. <code>b_index</code> is the rank of b in the aggregate tile of data. </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[out]</td><td class="paramname">tail_flags</td><td>Calling thread's discontinuity tail_flags </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">input</td><td>Calling thread's input items </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">flag_op</td><td>Binary boolean flag predicate </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">tile_successor_item</td><td><b>[<em>thread</em><sub><code>BLOCK_THREADS</code>-1</sub> only]</b> Item with which to compare the last tile item (<code>input</code><sub><em>ITEMS_PER_THREAD</em>-1</sub> from <em>thread</em><sub><em>BLOCK_THREADS</em>-1</sub>). </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__discontinuity_8cuh_source.html#l00547">547</a> of file <a class="el" href="block__discontinuity_8cuh_source.html">block_discontinuity.cuh</a>.</p>

</div>
</div>
<hr/>The documentation for this class was generated from the following file:<ul>
<li><a class="el" href="block__discontinuity_8cuh_source.html">block_discontinuity.cuh</a></li>
</ul>
</div><!-- contents -->
<!-- HTML footer for doxygen 1.8.3.1-->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated on Fri Aug 23 2013 17:31:11 for CUB by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.4
<br>
&copy; 2013 NVIDIA Corporation
</small></address>
</body>
</html>
